#pragma once

#if defined(__cplusplus) && defined(__CUDACC__)

#if defined(__CUDACC_RTC__)
#define __DEVICE_FUNCTIONS_DECL__ __device__
#define __DEVICE_FUNCTIONS_STATIC_DECL__ __device__
#else
#define __DEVICE_FUNCTIONS_DECL__ __device__
#define __DEVICE_FUNCTIONS_STATIC_DECL__ static __inline__ __device__
#endif /* __CUDACC_RTC__ */

#include "builtin_types.h"
#include "device_types.h"
#include "host_defines.h"

__DEVICE_FUNCTIONS_STATIC_DECL__ int mulhi(const int a, const int b)
{
  return __mulhi(a, b);
}

__DEVICE_FUNCTIONS_STATIC_DECL__ unsigned int mulhi(const unsigned int a, const unsigned int b)
{
  return __umulhi(a, b);
}

__DEVICE_FUNCTIONS_STATIC_DECL__ unsigned int mulhi(const int a, const unsigned int b)
{
  return __umulhi(static_cast<unsigned int>(a), b);
}

__DEVICE_FUNCTIONS_STATIC_DECL__ unsigned int mulhi(const unsigned int a, const int b)
{
  return __umulhi(a, static_cast<unsigned int>(b));
}

__DEVICE_FUNCTIONS_STATIC_DECL__ long long int mul64hi(const long long int a, const long long int b)
{
  return __mul64hi(a, b);
}

__DEVICE_FUNCTIONS_STATIC_DECL__ unsigned long long int mul64hi(const unsigned long long int a, const unsigned long long int b)
{
  return __umul64hi(a, b);
}

__DEVICE_FUNCTIONS_STATIC_DECL__ unsigned long long int mul64hi(const long long int a, const unsigned long long int b)
{
  return __umul64hi(static_cast<unsigned long long int>(a), b);
}

__DEVICE_FUNCTIONS_STATIC_DECL__ unsigned long long int mul64hi(const unsigned long long int a, const long long int b)
{
  return __umul64hi(a, static_cast<unsigned long long int>(b));
}

__DEVICE_FUNCTIONS_STATIC_DECL__ int float_as_int(const float a)
{
  return __float_as_int(a);
}

__DEVICE_FUNCTIONS_STATIC_DECL__ float int_as_float(const int a)
{
  return __int_as_float(a);
}

__DEVICE_FUNCTIONS_STATIC_DECL__ unsigned int float_as_uint(const float a)
{
  return __float_as_uint(a);
}

__DEVICE_FUNCTIONS_STATIC_DECL__ float uint_as_float(const unsigned int a)
{
  return __uint_as_float(a);
}
__DEVICE_FUNCTIONS_STATIC_DECL__ float saturate(const float a)
{
  return __saturatef(a);
}

__DEVICE_FUNCTIONS_STATIC_DECL__ int mul24(const int a, const int b)
{
  return __mul24(a, b);
}

__DEVICE_FUNCTIONS_STATIC_DECL__ unsigned int umul24(const unsigned int a, const unsigned int b)
{
  return __umul24(a, b);
}

__DEVICE_FUNCTIONS_STATIC_DECL__ int float2int(const float a, const enum cudaRoundMode mode)
{
  return (mode == cudaRoundNearest) ? __float2int_rn(a) : (mode == cudaRoundPosInf) ? __float2int_ru(a)
                                                      : (mode == cudaRoundMinInf)   ? __float2int_rd(a)
                                                                                    : __float2int_rz(a);
}

__DEVICE_FUNCTIONS_STATIC_DECL__ unsigned int float2uint(const float a, const enum cudaRoundMode mode)
{
  return (mode == cudaRoundNearest) ? __float2uint_rn(a) : (mode == cudaRoundPosInf) ? __float2uint_ru(a)
                                                       : (mode == cudaRoundMinInf)   ? __float2uint_rd(a)
                                                                                     : __float2uint_rz(a);
}

__DEVICE_FUNCTIONS_STATIC_DECL__ float int2float(const int a, const enum cudaRoundMode mode)
{
  return (mode == cudaRoundZero) ? __int2float_rz(a) : (mode == cudaRoundPosInf) ? __int2float_ru(a)
                                                   : (mode == cudaRoundMinInf)   ? __int2float_rd(a)
                                                                                 : __int2float_rn(a);
}

__DEVICE_FUNCTIONS_STATIC_DECL__ float uint2float(const unsigned int a, const enum cudaRoundMode mode)
{
  return (mode == cudaRoundZero) ? __uint2float_rz(a) : (mode == cudaRoundPosInf) ? __uint2float_ru(a)
                                                    : (mode == cudaRoundMinInf)   ? __uint2float_rd(a)
                                                                                  : __uint2float_rn(a);
}

#undef __DEVICE_FUNCTIONS_DECL__
#undef __DEVICE_FUNCTIONS_STATIC_DECL__

#endif /* __cplusplus && __CUDACC__ */
